[Cherry-Pick][Operator Mechanism] fix Broadcast support for big tensor(#79439) by feixi139 · Pull Request #79440 · PaddlePaddle/Paddle

feixi139 · 2026-07-08T07:42:15Z

PR Category

Operator Mechanism

PR Types

Bug fixes

Description

BroadcastConfig 中使用 FastDivMod 对输出线性索引进行维度坐标拆解。当 tensor 某一维度大小超过 INT_MAX（2^31 - 1）时，int 类型发生溢出，导致索引计算错误，广播结果产生静默错误。

修复

将 BroadcastConfig 中的 FastDivMod 替换为 FastDivMod<int64_t>，移除 PADDLE_ENFORCE_LE_INT_MAX 的限制，使广播操作支持单维度超过 2^31 的大 tensor。

devPR:#79439

是否引起精度变化

否

fix bugs

PaddlePaddle-bot · 2026-07-09T09:36:52Z

🤖 Paddle-CI-Agent | ci_status_monitor | 2026-07-15 17:02:53 UTC+08:00

CI报告基于以下代码生成（30分钟更新一次）:
PR commit: e8ece39 | Merge base: eaa369b (branch: release/3.4)

1 Required任务 : 22/32 通过

总执行（rerun次数）	总任务	✅ 通过	❌ 失败	⏳ 运行中	⏸️ 等待中	跳过
57(0)	57	46	2	9	0	0

任务	错误类型	置信度	日志
`Check approval`	需要 Approval	高	Job

2 失败详情

🔴 Check approval — 需要 Approval（置信度: 高）

该 Job 需要人工 Approval，完成审批后 CI 才会继续执行。

修复建议：请通过人工审批。

paddle-bot · 2026-07-14T14:00:11Z

你的PR提交成功，感谢你对开源项目的贡献!
请关注后续CI自动化测试结果，详情请参考Paddle-CI手册。
Your PR has been submitted. Thanks for your contribution!
Please wait for the result of CI firstly. See Paddle CI Manual for details.

PaddlePaddle-bot

🤖 Paddle-CI-Agent | pr_review | 2026-07-15 15:59:03

📋 Review 摘要

PR 概述：修复 BroadcastConfig 在大 tensor 广播索引拆解中的 32 位溢出问题。
变更范围：paddle/phi/kernels/funcs/broadcast_function.h、paddle/phi/kernels/primitive/datamover_primitives.h
影响面 Tag：[Operator Mechanism]

问题

未发现阻塞性问题。PR 规范问题在下面章节报，不在这里重复。

📝 PR 规范检查

标题不符合 release 分支 cherry-pick 格式：当前 PR body 已包含 develop 原 PR 链接，但标题没有使用 checklist 要求的 [Cherry-Pick][Tag] 标题描述(#原PR号) 格式。描述结构符合规范。

标题建议（可直接复制）：

[Cherry-Pick][Operator Mechanism] fix Broadcast support for big tensor(#79439)

PR 描述建议（点击展开，可直接复制）

### PR Category
<!-- One of [ User Experience | Execute Infrastructure | Operator Mechanism | CINN | Custom Device | Performance Optimization | Distributed Strategy | Parameter Server | Communication Library | Auto Parallel | Inference | Environment Adaptation ] -->
Operator Mechanism

### PR Types
<!-- One of [ New features | Bug fixes | Improvements | Performance | BC Breaking | Deprecations | Docs | Devs | Not User Facing | Security | Others ] -->
Bug fixes

### Description
<!-- Describe what you’ve done -->
BroadcastConfig 中使用 FastDivMod<int> 对输出线性索引进行维度坐标拆解。当 tensor 某一维度大小超过 INT_MAX（2^31 - 1）时，int 类型发生溢出，导致索引计算错误，广播结果产生静默错误。

本 PR 将 BroadcastConfig 中的 FastDivMod<int> 替换为 FastDivMod<int64_t>，并将相关广播索引和 stride 累乘提升到 int64_t，以支持单维度超过 2^31 的大 tensor。

devPR:https://github.com/PaddlePaddle/Paddle/pull/79439

### 是否引起精度变化
<!-- one of the following [ 是 | 否 ]-->
否

总体评价

代码层面围绕 big tensor 的 divmod、stride、block offset 与 ReadDataBc 调用链做了回溯，未发现新的阻塞性问题。XPU 路径仍依赖 32 位 datamover，并在构造 BroadcastConfig 前保留 INT_MAX 校验，符合当前实现约束。

A-nnonymous · 2026-07-15T09:43:11Z

 */
 struct BroadcastConfig {
-  funcs::FastDivMod<int> divmoders[DDim::kMaxRank];
+  funcs::FastDivMod<int64_t> divmoders[DDim::kMaxRank];


这里当前无条件使用 FastDivMod<int64_t>，会让普通小 tensor 广播也走 64 位 divmod；但本 PR 的 launcher 仍专门选择 uint32_t 快路径。源 PR #79439 最新 head 已将 BroadcastConfig 模板化为 BroadcastConfig，小 tensor 保留 32 位、仅大 tensor 用 64 位。建议同步该最新修正，避免 release 分支引入全量 broadcast 热路径的性能回退。

wanghuancoder · 2026-07-15T09:43:49Z

  static __device__ __forceinline__ void Apply(const Array &ins,
                                               ArgsT *args,
-                                               uint32_t index_bc[][VecSize]) {
+                                               uint64_t index_bc[][VecSize]) {


这里可以用IndexT吗？

wanghuancoder · 2026-07-15T09:44:40Z

 #else
  if (LoadType == kBroadcast) {
-    uint32_t index_bc[Arity][VecSize] = {0};
+    uint64_t index_bc[Arity][VecSize] = {0};


这里可以用IndexT吗？

wanghuancoder · 2026-07-15T09:44:52Z

 #pragma unroll
    for (int k = 0; k < VecSize; ++k) {
-      uint32_t idx = thread_offset + k;
+      uint64_t idx = thread_offset + k;


这里可以用IndexT吗？

wanghuancoder · 2026-07-15T09:45:16Z

+  uint64_t block_offset =
+      static_cast<uint64_t>(BLOCK_ID_X) * BLOCK_NUM_X * read_lens;
+  uint64_t stride = static_cast<uint64_t>(BLOCK_NUM_X) * GRID_NUM_X * read_lens;
+  for (; block_offset < static_cast<uint64_t>(main_offset);
+       block_offset += stride) {


这里可以用IndexT吗？

wanghuancoder · 2026-07-15T09:45:30Z

+  int64_t num =
+      static_cast<int64_t>(numel) - static_cast<int64_t>(block_offset);


这里可以用IndexT吗？

wanghuancoder · 2026-07-15T09:46:14Z

                                            configs,
-                                            tail_tid,
-                                            block_offset,
+                                            static_cast<uint32_t>(num),


这里为啥用int32？

wanghuancoder · 2026-07-15T09:47:19Z

                              NumOuts,
                              VecSize,
-                              kElementwise>
+                              false>


这里为啥要改成false？

我知道了，你还是需要XPU和GPU分开

wanghuancoder · 2026-07-15T09:48:56Z

+    const IndexT numel = static_cast<IndexT>(classifier.numel);
+    const int threads = 64;
+    const int blocks = 8;
+    int read_lens = classifier.configs[0].buf_len;
+    auto stream = dev_ctx.x_context()->xpu_stream;
+    const IndexT block_len = static_cast<IndexT>(read_lens) * threads;
+    const IndexT main_offset = (numel / block_len) * block_len;
+    const IndexT tail_tid = numel % block_len;


这里在GPU Kernel外部，可以用int64

wanghuancoder · 2026-07-15T09:50:21Z

 struct BroadcastConfig {
-  funcs::FastDivMod<int> divmoders[DDim::kMaxRank];
+  funcs::FastDivMod<int64_t> divmoders[DDim::kMaxRank];


这个Config在Broadcast Kernel里使用。因此是GPU运算使用的代码。需要用IndexT区分

wanghuancoder · 2026-07-15T09:51:57Z

+  launch_with_index_t(uint32_t{0});
+#else
+  const uint64_t index_limit = std::numeric_limits<uint32_t>::max();
+  const uint64_t grid_stride =
+      static_cast<uint64_t>(gpu_config.block_per_grid.x) *
+      static_cast<uint64_t>(gpu_config.GetBlockSize()) * VecSize;
+  if (static_cast<uint64_t>(classifier.numel) <= index_limit &&
+      grid_stride <= index_limit) {
+    launch_with_index_t(uint32_t{0});
  } else {
-    VectorizedBroadcastKernel<Functor, OutT, Arity, NumOuts, VecSize, kMixed>
-        <<<blocks, threads, 0, stream>>>(classifier.ins_data,
-                                         classifier.outs_data,
-                                         classifier.use_broadcast,
-                                         numel,
-                                         classifier.configs,
-                                         main_offset,
-                                         tail_tid,
-                                         VecSize,
-                                         func);
+    launch_with_index_t(uint64_t{0});


XPU的逻辑可以不用动。
GPU的逻辑区分std::numeric_limits<uint32_t>::max()即可。
结合实际情况看看launch_with_index_t是否必要，如果必要就封装，不必要可以平铺写。

fix bugs

b0efc76

fix bugs

feixi139 changed the title ~~[release/3.4][Operator Mechanism] fix Broadcast do not support big_tensor~~ [release/3.4][Operator Mechanism] fix BroadcastConfig do not support big_tensor Jul 8, 2026